/**
 * Copyright (c) 2009, Robert Schuster
 * All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or
 * without modification, are permitted provided that the
 * following conditions are met:
 *
 * - Redistributions of source code must retain the above
 *   copyright notice, this list of conditions and the
 *   following disclaimer.
 * - Redistributions in binary form must reproduce the above
 *   copyright notice, this list of conditions and the
 *   following disclaimer in the documentation and/or other
 *   materials provided with the distribution.
 * - Neither the name of Robert Schuster nor the names of his
 *   contributors may be used to endorse or promote products
 *   derived from this software without specific prior written
 *   permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND
 * CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
 * INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF
 * MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR
 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
 * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT
 * NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
 * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
 * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
 * OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

/* 
 * File:   Executor.cpp
 * Author: Robert Schuster
 * 
 * Created on 3. September 2009, 14:03
 */

#include "Executor.hpp"
#include "CLcalc.h"
#include <vector>
#include <boost/tokenizer.hpp>

namespace CLcalc {

Executor::Executor(Device* device) {
    //add the device
    devices.push_back(device);
    Context = NULL;
    Queue = NULL;
}

Executor::Executor(const Executor& orig) {
    cout << "ERROR: Never copy an Executor-Object!" << endl;
}

Executor::~Executor() {
    //cout << "Executor::~Executor()" << endl;
    //Release the Memory Objects
    for (int i=0;i<Memorys.size();i++) {
        releaseData(i);
    }

    //Release the Devices, the Environment takes care to delete them
    for (int i=0;i<devices.size();i++) {
        //wait for everything on the device to finish
        if (Queue != NULL) waitForAll(i);
        devices.release(devices.begin()).release();
    }

    //Release the Queue and the Context
    if (Context != NULL) {
        //Release the Command Queue
        for (int i=0;i<devices.size();i++) {
            clReleaseCommandQueue(Queue[i]);
        }

        //Release the Context for this executor
        clReleaseContext(*Context);

        //Delete the Context
        delete Context;
    }
    //cout << "Executor::~Executor()" << endl;
}

int Executor::addProgFromSource(const char* source, const char* name) throw (Error) {
    //if the context is null, init it!
    if (Context == NULL) Init_Context();

    //check if this Program is already compiled on all devices
    int found = -1;
    for (int i=0;i<Programs.size();i++) {
        //compare the name of the Program
        if (string(Programs[i].getName()).compare(name) == 0) {
            found = i;
            break;
        }
    }

    //not compiled? then do it!
    if (found == -1) {
        //run the preprocessor for templates
        string strsource(source);
        string strname(name);
        Preprocessor(strsource, strname);
        //compile it
        CompiledProgram* cp = new CompiledProgram(strsource.c_str(), name , Context);
        Programs.push_back(cp);
        //return the index in the vector as ID
        return Programs.size()-1;
    } else {
        return found;
    }

}

void Executor::Init_Context() throw (Error) {
    if (Context != NULL) throw Error(-1);

    //create a c-Array of the device-ids
    cl_device_id ids[devices.size()];
    for (int i=0;i<devices.size();i++) {
        ids[i] = devices[i].getID();
    }
    //Create a Context for this Devices
    int err = CL_SUCCESS;
    Context = new cl_context;
    *Context = clCreateContext(0, devices.size(), &ids[0], NULL, NULL, &err);
    if (err != CL_SUCCESS) {
        delete Context;
        throw Error(err);
    }

    //Create a Command Queue for each Device
    Queue = new cl_command_queue[devices.size()];
    for (int i=0;i<devices.size();i++) {
        Queue[i] = clCreateCommandQueue(*Context, ids[0], 0, &err);
        if (err != CL_SUCCESS) {
            throw Error(err);
        }        
    }

}

int Executor::addData(void* data, size_t data_size, int data_dir) throw (Error) {
    int DeviceID = 0;
    //if the context is null, init it!
    if (Context == NULL) Init_Context();

    //Create the cl_mem object
    cl_mem* cldata = new cl_mem;
    int err = CL_SUCCESS;
    if ((data_dir == CL_MEM_USE_HOST_PTR) || (data_dir == CL_MEM_COPY_HOST_PTR))
        *cldata = clCreateBuffer(*Context, data_dir, data_size, data, &err);
    else
        *cldata = clCreateBuffer(*Context, data_dir, data_size, NULL, &err);
    if (err != CL_SUCCESS)
    {
        delete cldata;
        throw Error(err);
    }

    // Write our data set into the input array in device memory
    if ((data_dir == CL_MEM_READ_ONLY) || (data_dir == CL_MEM_READ_WRITE))
        err = clEnqueueWriteBuffer(Queue[DeviceID], *cldata, CL_TRUE, 0, data_size, data, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        throw Error(err);
    }
    
    //Save the pointer and return its index
    Memorys.push_back(cldata);
    return Memorys.size()-1;
}

void Executor::getData(int DataID, void* data, size_t data_size) throw (Error) {
    int DeviceID = 0;
    //is the DataID valid?
    if (Memorys.size() <= DataID) throw Error(-1);

    //if the size if not known, ask the memory object
    int err;
    if (data_size == -1) {
        err = clGetMemObjectInfo(Memorys[DataID], CL_MEM_SIZE, sizeof(size_t), &data_size, NULL);
        if (err != CL_SUCCESS) throw Error(err);
    }

    //read the Data
    err = clEnqueueReadBuffer(Queue[DeviceID], Memorys[DataID], CL_TRUE, 0, data_size, data, 0, NULL, NULL );
    if (err != CL_SUCCESS) throw Error(err);
}

void Executor::releaseData(int DataID) throw (Error) {
    //is the DataID valid?
    if (Memorys.size() <= DataID) throw Error(-1);
    //is the object != null? then release it and remove it from the vector
    if (!Memorys.is_null(DataID)) {
        int err = clReleaseMemObject(Memorys[DataID]);
        if (err != CL_SUCCESS) throw Error(err);
        Memorys.replace(DataID, NULL);
    }
}

void* Executor::getDataHostPtr(int DataID) throw (Error) {
    //is the DataID valid?
    if (Memorys.size() <= DataID) throw Error(-1);
    //is the object != null? then release it and remove it from the vector
    if (!Memorys.is_null(DataID)) {
        void* ptr = NULL;
        clGetMemObjectInfo(Memorys[DataID], CL_MEM_HOST_PTR, sizeof(void*), &ptr, NULL);
        return ptr;
    } else {
        return NULL;
    }
}

long Executor::getDataDir(int DataID) throw (Error) {
    //is the DataID valid?
    if (Memorys.size() <= DataID) throw Error(-1);
    //is the object != null? then release it and remove it from the vector
    if (!Memorys.is_null(DataID)) {
        cl_mem_flags dir = -1;
        clGetMemObjectInfo(Memorys[DataID], CL_MEM_FLAGS, sizeof(cl_map_flags), &dir, NULL);
        return dir;
    } else {
        return -1;
    }

}


void Executor::setArgument(int progID, int arg_pos , int DataID) throw (Error) {
    //is the progID valid?
    if (Programs.size() <= progID) throw Error(-1);
    //is the dataID valid?
    if (Memorys.size() <= DataID) throw Error(-1);

    //set the argument to the kernel
    int err  = clSetKernelArg(*(Programs[progID].getKernel()), arg_pos, sizeof(cl_mem), &Memorys[DataID]);
    if (err != CL_SUCCESS) throw Error(err);
}

void Executor::setConstIntArgument(int progID, int arg_pos, int argument) throw (Error) {
    //is the progID valid?
    if (Programs.size() <= progID) throw Error(-1);
    //set the argument to the kernel
    int err  = clSetKernelArg(*(Programs[progID].getKernel()), arg_pos, sizeof(int), &argument);
    if (err != CL_SUCCESS) throw Error(err);
}

void Executor::Execute(int progID, int MaxIndex) throw (Error) {
    int DeviceID = 0;
    //is the progID valid?
    if (Programs.size() <= progID) throw Error(-1);
    //Execute the kernel
    size_t range = MaxIndex;
    int err = clEnqueueNDRangeKernel(Queue[DeviceID], *(Programs[progID].getKernel()), 1, NULL, &range, NULL, 0, NULL, NULL);
    waitForAll(DeviceID);
    if (err != CL_SUCCESS) throw Error(err);
}

void Executor::Execute(int progID, int MaxIndex, int wg_size) throw (Error) {
    int DeviceID = 0;
    //is the progID valid?
    if (Programs.size() <= progID) throw Error(-1);
    //Execute the kernel
    size_t range = MaxIndex;
    size_t local_work_size = wg_size;
    int err = clEnqueueNDRangeKernel(Queue[DeviceID], *(Programs[progID].getKernel()), 1, NULL, &range, &local_work_size, 0, NULL, NULL);
    waitForAll(DeviceID);
    if (err != CL_SUCCESS) throw Error(err);
}

Device* Executor::getDevice(int DeviceID) throw (Error) {
    //is the DeviceID valid?
    if (devices.size() <= DeviceID) throw Error(-1);
    return &devices[DeviceID];
}

void Executor::waitForAll(int DeviceID) throw (Error) {
    //is the DeviceID valid?
    if (devices.size() <= DeviceID) throw Error(-1);
    //is there a Queue for this device
    if (Queue == NULL) throw Error(-2);
    //wait for everything to finish
    int err = clFinish(Queue[DeviceID]);
    if (err != CL_SUCCESS) throw Error(err);
}

void Executor::Preprocessor(string &source, string &name) throw (Error) {
    //are there templates inside the source?
    size_t template_index = source.find("template");
    if (template_index != string::npos) {
        //we found a template, so we need a tokenizer
        typedef boost::tokenizer<boost::char_separator<char> > tokenizer;
        //find the first typename
        size_t temp = source.find("typename");
        if (temp != string::npos) {
            //now we are here:
            //template <typename T>
            //          ^
            //go behind          ^
            size_t behind = source.find(">", temp);
            if (behind == string::npos) {
                Error err(-1);
                err.Text = "missing < > inside the template definition! Sample: template <typename T>";
                throw err;
            }
            //create a string with the typeid
            string templateline = source.substr(temp, behind-temp);
            //create a tokenizer for the template line
            boost::char_separator<char> sep(" ,");
            tokenizer typenameTemplates(templateline, sep);
            //find the type inside the kernelname
            size_t tname = name.find("<");
            if (tname != string::npos) {
                tname++;
                size_t tname2 = name.find(">");
                if (tname2 == string::npos) {
                    Error err(-1);
                    err.Text = "missing < > inside the kernelname! Sample: sqrt<float>";
                    throw err;
                }
                //create a tokenizer
                tokenizer typenameNames(name.substr(tname,tname2-tname), sep);
                //create a string with the source without the template definition
                string newsource = source.erase(template_index, behind-template_index+1);
                //replace the typeid inside the kernel
                size_t pos = -1;
                //loop over all typenames
                vector<string> typenameNamesVector;
                size_t nullpos;
                for (tokenizer::iterator tok_n = typenameNames.begin(); tok_n != typenameNames.end(); ++tok_n) {
                    if ((nullpos = (*tok_n).find('\0')) != string::npos) {
                        string withoutnull(*tok_n);
                        withoutnull = withoutnull.erase(nullpos,1);
                        typenameNamesVector.push_back(withoutnull);
                    } else
                    typenameNamesVector.push_back(*tok_n);
                }
                vector<string>::iterator tok_names = typenameNamesVector.begin();
                //loop over all templates
                for (tokenizer::iterator tok_temp = typenameTemplates.begin(); tok_temp != typenameTemplates.end(); ++tok_temp) {
                    //is the token a string to replace?
                    if ((*tok_temp).compare("typename") == 0) continue;
                    //replace
                    while (string::npos != (pos = newsource.find(*tok_temp,pos+1))) {
                        //check the characters before ...
                        bool before_ok = false;
                        if (pos > 0) {
                            if (newsource.at(pos-1) == ' ') {
                                before_ok = true;
                            } else if (newsource.at(pos-1) == '\n') {
                                before_ok = true;
                            } else if (newsource.at(pos-1) == '\r') {
                                before_ok = true;
                            } else if (newsource.at(pos-1) == '\t') {
                                before_ok = true;
                            } else if (newsource.at(pos-1) == '(') {
                                before_ok = true;
                            }
                        } else {
                            before_ok = true;
                        }
                        //check the characters behind ...
                        bool behind_ok = false;
                        if (pos < newsource.length()-(*tok_temp).length()-2) {
                            if (newsource.at(pos+(*tok_temp).length()) == ' ') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == '\n') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == '\r') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == '\t') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == '*') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == ')') {
                                behind_ok = true;
                            } else if (newsource.at(pos+(*tok_temp).length()) == '(') {
                                behind_ok = true;
                            }
                        }
                        //replace if before and behind is ok
                        if (before_ok && behind_ok)
                            newsource = newsource.replace(pos,(*tok_temp).length(),*tok_names);
                    } // while loop
                    //go to the next name
                    if (tok_names != typenameNamesVector.end()) tok_names++;
                } //for loop
                source = newsource;
            } else {
                Error err(-1);
                err.Text = "no typename inside the kernelname! Sample: sqrt<float>";
                throw err;
            }            
        }
    }
}

//find the size of the datatype
size_t Executor::getDatatypSizeFromName(const char* templatetype) {
    size_t datasize = -1;
    if (strcmp("bool", templatetype) == 0) {
        datasize = 1;
    } else if (strcmp("char", templatetype) == 0) {
        datasize = 1;
    } else if (strcmp("short", templatetype) == 0) {
        datasize = 2;
    } else if (strcmp("int", templatetype) == 0) {
        datasize = 4;
    } else if (strcmp("long", templatetype) == 0) {
        datasize = 8;
    } else if (strcmp("float", templatetype) == 0) {
        datasize = 4;
    } else if (strcmp("double", templatetype) == 0) {
        datasize = 8;
    }
    return datasize;
}

/**
 * The Buildin Functions-----------------------------------------------------------------
 */

        //Kernels that are usefull for more then one function are here listed:
        namespace BuildIn {

            extern const char* kernel_buildin_basic1d;

            extern const char* kernel_buildin_CLintrinsics_1arg;
        }

        void Executor::BuildIn_basic(void* arg1, void* arg2, void* result, int* length, int dims, const char* fname, const char* tname) throw (Error) {
            int arg1ID = -1;
            int arg2ID = -1;
            int resultID = -1;
            Error* err = NULL;

            try {
                //calculate the size of the complete array
                size_t size = length[0];
                for (int i=1;i<dims; i++) {
                    size *= length[i];
                }

                //copy the data to the device and allocate the result array on the device
                size_t elementsize = getDatatypSizeFromName(tname);
                arg1ID = addData(arg1, elementsize*size, CL_MEM_USE_HOST_PTR);
                arg2ID = addData(arg2, elementsize*size, CL_MEM_USE_HOST_PTR);
                resultID = addData(result, elementsize*size, CL_MEM_USE_HOST_PTR);

                //execute the kernel
                BuildIn_basic_execOnly(arg1ID, arg2ID, resultID, length, dims, fname, tname);

            } catch (Error er) {
                err = new Error(er);
            }

            //cleanup the memory
            if (arg1ID != -1) releaseData(arg1ID);
            if (arg2ID != -1) releaseData(arg2ID);
            if (resultID != -1) releaseData(resultID);

            //throw an error?
            if (err != NULL) throw *err;
        }

        void Executor::BuildIn_basic_execOnly(int sum1ID, int sum2ID, int resultID, int* length, int dims, const char* fname, const char* tname) throw (Error) {
            int kernelid;

            //create the kernelname
            string kernelname("buildin_basic1d<");
            kernelname.append(tname);
            kernelname.append(",");
            kernelname.append(fname);
            kernelname.append(">");

            //compile the kernel
            kernelid = addProgFromSource(BuildIn::kernel_buildin_basic1d, kernelname.c_str());

            //calculate the size of the complete array
            int size = length[0];
            for (int i=1;i<dims; i++) {
                size *= length[i];
            }

            //set the kernel arguments
            setArgument(kernelid, 0, sum1ID);
            setArgument(kernelid, 1, sum2ID);
            setArgument(kernelid, 2, resultID);

            //execute the kernel
            Execute(kernelid, size);
        }

        void Executor::BuildIn_CLintrinsics1arg(void* input, void* output, int* length, int dims, const char* fname, const char* tname) throw (Error) {
            int arg1ID = -1;
            int arg2ID = -1;
            Error* err = NULL;

            try {
                //calculate the size of the complete array
                size_t size = length[0];
                for (int i=1;i<dims; i++) {
                    size *= length[i];
                }

                //copy the data to the device and allocate the result array on the device
                size_t elementsize = getDatatypSizeFromName(tname);
                arg1ID = addData(input, elementsize*size, CL_MEM_USE_HOST_PTR);
                arg2ID = addData(output, elementsize*size, CL_MEM_USE_HOST_PTR);

                //Execute the Kernel
                BuildIn_CLintrinsics1arg_execOnly(arg1ID, arg2ID, length, dims, fname, tname);
            } catch (Error er) {
                err = new Error(er);
            }

            //cleanup the memory
            if (arg1ID != -1) releaseData(arg1ID);
            if (arg2ID != -1) releaseData(arg2ID);

            //throw an error?
            if (err != NULL) throw *err;
        }

        void Executor::BuildIn_CLintrinsics1arg_execOnly(int inputID, int outputID, int* length, int dims, const char* fname, const char* tname) throw (Error) {
            int progID = -1;
            try {
                //calculate the size of the complete array
                size_t size = length[0];
                for (int i=1;i<dims; i++) {
                    size *= length[i];
                }

                //create the kernelname
                string kernelname("buildin_CLintrinsics_1arg<");
                kernelname.append(tname);
                kernelname.append(",");
                kernelname.append(fname);
                kernelname.append(">");

                //compile the kernel
                progID = addProgFromSource(BuildIn::kernel_buildin_CLintrinsics_1arg, kernelname.c_str());

                //set the arguments
                setArgument(progID, 0, inputID);
                setArgument(progID, 1, outputID);

                //Execute the Kernel
                Execute(progID, size);
            } catch (Error er) {
                throw er;
            }
        }

} // end of namespace CLcalc